//-------------------------------------------------------------------
//	Copyright (c) 2013-2014  Zhirnov Andrey
//	This file is part of the "UX OpwnGL Framework" project.
//	See copyright notice in "UX_GL_Framework.h".
//-------------------------------------------------------------------

#include "ShaderLoader.h"

namespace Framework
{

	//
	// Shader Manager
	//

	class ShaderManager
	{
	// types
	private:
		typedef static_array< GLShaderPtr, gl_shader::_IDX_COUNT >	shader_st_array_t;

		struct GLShaderLib
		{
		// variables
		public:
			shader_st_array_t	shaders;
			string				header;
			string				source;


		// methods
		public:
			GLShaderLib ();

			bool Load (uni_c_string filename, uni_c_string preprocessor);

			GLShaderPtr Get (uni_c_string selfFileName, gl_shader::bits shaderBit);
		};

		typedef map< string, GLShaderLib >	glshader_libs_t;


	// variables
	private:
		// GL //
		glshader_libs_t	_glShaderLibs;
		string			_glslHeader;

		// CL //
		string					_clHeader;
		cl_prog_options::type	_clOptions;


	// methods
	public:
		ShaderManager ();
		
		bool Load (GLProgramPtr &program, uni_c_string filename, uni_c_string preprocessor);
		bool Load (CLProgramPtr &program, uni_c_string filename, uni_c_string preprocessor, uni_c_string options);

		void FreeCache ();

		static ShaderManager * Instance ()
		{
			return SingletonMultiThread::Instance< ShaderManager >();
		}


	private:
		bool _LoadLibrary (uni_c_string filename, OUT GLShaderLib *&lib);
		bool _ProcessSource (array<GLShaderPtr> &shaders, gl_shader::bits shaderBit, INOUT string &src, uni_c_string path);

		bool _Load (array<GLShaderPtr> &shaders, uni_c_string filename, uni_c_string preprocessor);

		static bool _ParseString (uni_c_string src, INOUT usize &pos, OUT uni_c_string &filename);

		static void _BuildGLSLProprocessor (OUT string &str);
		static void _BuildCLPreprocessor (OUT string &str, OUT cl_prog_options::type &opt);

		static bool _ReadFile (uni_c_string filename, OUT string &buf);
		static bool _RecursiveParseCLFile (uni_c_string filename, OUT string &src, INOUT set< FileAddress > &importedFiles);
	};
	

	
//-------------------------------- GLShaderLib ------------------------------//
/*
=================================================
	constructor
=================================================
*/
	ShaderManager::GLShaderLib::GLShaderLib ()
	{
	}
	
/*
=================================================
	Load
=================================================
*/
	bool ShaderManager::GLShaderLib::Load (uni_c_string filename, uni_c_string preprocessor)
	{
		RFilePtr	file;
		CHECK_VALUE( (file = File::HddRFile::New( filename )).IsNotNull() );
		
		const usize	len = file->Size();
		
		source.ClearMem();
		source.Reserve( len + 1 );

		CHECK_RETURN( file->Read( source.Begin(), len ) );
		source.SetLength( len );

		static const uni_c_string	s_header( "[HEADER]" );
		static const uni_c_string	s_lib( "[LIB]" );

		usize	header_begin	= 0,
				header_end		= 0,
				lib_begin		= 0;

		source.Find( s_header, header_begin, 0 );
		source.Find( s_lib, lib_begin, 0 );

		header_end = lib_begin;
		StringParser::ToNextLine( source, header_begin );
		StringParser::ToNextLine( source, lib_begin );
		StringParser::ToPrevLine( source, header_end );

		ASSERT( header_end > header_begin );

		header = uni_c_string( source.cstr() + header_begin, header_end - header_begin );
		
		const usize	header_first_line	= StringParser::CalculateNumberOfLines( uni_c_string( source.cstr(), header_begin ) );
		const usize	source_first_line	= StringParser::CalculateNumberOfLines( uni_c_string( source.cstr(), lib_begin ) );

		if ( lib_begin > 0 )
			source.Erase( 0, lib_begin );

		CHECK_VALUE( not header.Empty() );
		CHECK_VALUE( not source.Empty() );

		// bug with #line
		//header.Insert( string("#line ") << (header_first_line+1) << " \"" << FileAddressUtils::GetName( filename ) << "-header\" \n", 0 );

		source.Insert( string("#line ") << (source_first_line+1) << " \"" << FileAddressUtils::GetName( filename ) << " \" \n", 0 );

		source.Insert( header, 0 );
		source.Insert( uni_c_string("#define LIB\n\n"), 0 );
		source.Insert( preprocessor, 0 );

		return true;
	}
	
/*
=================================================
	FreeCache
=================================================
*/
	void ShaderManager::FreeCache ()
	{
		_glShaderLibs.ClearMem();
	}
	
/*
=================================================
	Get
=================================================
*/
	GLShaderPtr ShaderManager::GLShaderLib::Get (uni_c_string selfFileName, gl_shader::bits shaderBit)
	{
		if ( shaders[ shaderBit ].IsNull() )
		{
			GLShaderPtr & shader = shaders[ shaderBit ];
			shader = new GLShader();

			CHECK_RETURN_R( shader->Create( gl_shader::BitToType( shaderBit ) ), null );
			CHECK_RETURN_R( shader->SetSource( source.cstr() ), null );
		
			if ( not shader->Compile() )
			{
				string	log;
				log << "In file \"" << selfFileName << "\":\n";

				shader->GetCompilationLog( log );
				RETURN_ERROR_R( log, null );
			}
		}
		
		return shaders[ shaderBit ];
	}


	
//------------------------------ ShaderManager ------------------------------//
/*
=================================================
	constructor
=================================================
*/
	ShaderManager::ShaderManager ()
	{
		_glShaderLibs.Reserve( 100 );

		_BuildGLSLProprocessor( _glslHeader );
		_BuildCLPreprocessor( _clHeader, _clOptions );
	}
	
/*
=================================================
	_LoadLibrary
=================================================
*/
	bool ShaderManager::_LoadLibrary (uni_c_string filename, OUT GLShaderLib *&lib)
	{
		usize	index;
		
		if ( _glShaderLibs.FindIndex( filename, index ) )
		{
			lib = &_glShaderLibs[ index ].second;
			return true;
		}

		index = _glShaderLibs.AddOrReplace( filename, GLShaderLib() );
		
		lib = &_glShaderLibs[ index ].second;

		CHECK_RETURN( lib->Load( filename, _glslHeader ) );

		return true;
	}
	
/*
=================================================
	_ParseString
=================================================
*/
	bool ShaderManager::_ParseString (uni_c_string src, INOUT usize &pos, OUT uni_c_string &filename)
	{
		for (; pos < src.Length(); ++pos)
		{
			const char	c = src[ pos ];

			if ( c == ' ' or c == '\t' )
				continue;

			if ( c == '"' )
				break;

			RETURN_ERROR( "incorrect char '" << c << "'" );
		}

		const usize	begin = ++pos;

		for (; pos < src.Length(); ++pos)
		{
			const char	c = src[ pos ];

			if ( c == '"' )
				break;

			if ( c == 10 or c == 13 ) {
				RETURN_ERROR( "not pair for braket \" " );
			}
		}

		filename = uni_c_string( src.cstr() + begin, pos - begin );
		return true;
	}

/*
=================================================
	_ProcessSource
----
	Warning: #impot work in commented code too!
=================================================
*/
	bool ShaderManager::_ProcessSource (array<GLShaderPtr> &shaders, gl_shader::bits shaderBit, INOUT string &src, uni_c_string path)
	{
		static const uni_c_string	s_import( "#import" );

		for (usize pos = 0; src.Find( s_import, pos, pos ); )
		{
			const usize	begin = pos;

			pos += s_import.Length();

			uni_c_string	import_lib;
			CHECK_RETURN( _ParseString( src, pos, import_lib ) );

			StringParser::ToNextLine( src, pos );

			GLShaderLib * lib = null;
			string		fname;
			FileAddressUtils::BuildPath( fname, path, import_lib, "slib" );

			CHECK_RETURN( _LoadLibrary( fname, lib ) );
			
			src.Erase( begin, pos - begin );
			src.Insert( lib->header, begin );

			pos  = begin + lib->header.Length();

			GLShaderPtr shader = lib->Get( fname, shaderBit );
			CHECK_VALUE( shader.IsNotNull() );

			shaders.PushBack( shader );
		}

		return true;
	}
	
/*
=================================================
	_Load
=================================================
*/
	bool ShaderManager::_Load (array<GLShaderPtr> &shaders, uni_c_string filename, uni_c_string preprocessor)
	{
		struct TConfigShaderType {
			uni_c_string	name;
			gl_shader::type	type;
		};

		const gl_shader::type	s_eSharedBlock	= gl_shader::type( gl_shader::_UNKNOWN - 1 );

		const TConfigShaderType	s_aShaderTypes[] = {{ "[VERTEX]",			gl_shader::VERTEX },
													{ "[TESS_CONTROL]",		gl_shader::TESS_CONTROL },
													{ "[TESS_EVALUATION]",	gl_shader::TESS_EVALUATION },
													{ "[GEOMETRY]",			gl_shader::GEOMETRY },
													{ "[FRAGMENT]",			gl_shader::FRAGMENT },
													{ "[COMPUTE]",			gl_shader::COMPUTE },
													{ "[END]",				gl_shader::_UNKNOWN },
													{ "[SHARED]",			s_eSharedBlock } };

		STATIC_ASSERTE( shader_st_array_t::STATIC_COUNT == COUNT_OF(s_aShaderTypes)-2 );

		
		RFilePtr	file;
		CHECK_VALUE( (file = File::HddRFile::New( filename )).IsNotNull() );
		
		string	s_source;
		usize	u_length = file->Size();
		usize	a_offsets[ COUNT_OF( s_aShaderTypes ) ] = {0};
		string	s_buf;

		s_source.Reserve( u_length + 1 );
		CHECK_RETURN( file->Read( s_source.ptr(), u_length ) );
		s_source.SetLength( u_length );
		

		for (usize i = 0; i < CountOf(s_aShaderTypes); ++i)
		{
			a_offsets[i] = usize(-1);

			if ( s_source.Find( s_aShaderTypes[i].name, a_offsets[i], 0 ) )
			{
				s_source[ a_offsets[i] ] = '\0';
				s_source[ (a_offsets[i] += s_aShaderTypes[i].name.Length() ) ] = '\0';
			}
		}

		uni_c_string	shared_src;

		for (usize i = 0; i < CountOf(s_aShaderTypes); ++i)
		{
			if ( s_aShaderTypes[i].type == s_eSharedBlock )
			{
				if ( a_offsets[i] == usize(-1) )
					break;
				
				shared_src = uni_c_string( s_source.cstr() + a_offsets[i]+1 );
				break;
			}
		}

		for (usize i = 0; i < CountOf(s_aShaderTypes); ++i)
		{
			if ( a_offsets[i] != usize(-1) and
				 s_aShaderTypes[i].type != gl_shader::_UNKNOWN and
				 s_aShaderTypes[i].type != s_eSharedBlock )
			{
				GLShaderPtr	p_shader = new GLShader();

				const usize	num_lines = StringParser::CalculateNumberOfLines( uni_c_string( s_source.cstr(), a_offsets[i]+1 ) );

				s_buf.ClearMem();
				s_buf	<< _glslHeader << '\n' << preprocessor << '\n' << shared_src << '\n'
						<< "#line " << (num_lines) << " \"" << FileAddressUtils::GetName( filename ) << " \" \n"
						<< uni_c_string( s_source.cstr() + a_offsets[i]+1 );

				CHECK_RETURN( _ProcessSource( shaders, gl_shader::TypeToBit( s_aShaderTypes[i].type ),
											  s_buf, FileAddressUtils::GetPath( filename ) ) );

				CHECK_RETURN( p_shader->Create( s_aShaderTypes[i].type ) );
				CHECK_RETURN( p_shader->SetSource( s_buf.cstr(), s_buf.Length() ) );
				
				if ( not p_shader->Compile() )
				{
					string	log;
					log << "In file \"" << filename << "\":\n";

					p_shader->GetCompilationLog( log );
					RETURN_ERROR( log );
				}

				shaders.PushBack( p_shader );
			}
		}

		return true;
	}

/*
=================================================
	Load
=================================================
*/
	bool ShaderManager::Load (GLProgramPtr &program, uni_c_string filename, uni_c_string preprocessor)
	{
		if ( program.IsNull() )
			program = new GLProgram();
		else
			program->Destroy();

		array< GLShaderPtr >	a_shaders;

		CHECK_RETURN( program->Create() );
		CHECK_RETURN( _Load( a_shaders, filename, preprocessor ) );

		FOR( i, a_shaders )
		{
			if ( a_shaders[i].IsNotNull() )
				CHECK_RETURN( program->Attach( a_shaders[i].ptr() ) );
		}
		
		if ( not program->Link() )
		{
			string	log;
			log << "In file \"" << filename << "\":\n";

			program->GetLinkingLog( log );
			RETURN_ERROR( log );
		}
		return true;
	}
	
/*
=================================================
	Load
=================================================
*/
	bool ShaderManager::Load (CLProgramPtr &program, uni_c_string filename, uni_c_string preprocessor, uni_c_string options)
	{
		if ( program.IsNull() )
			program = new CLProgram();
		else
			program->Destroy();
		
		string	src;
		src.Reserve( 1<<12 );	// 4kb

		CHECK_RETURN( _ReadFile( filename, src ) );
		
		const uni_c_string		path = FileAddressUtils::GetPath( filename );

		set< FileAddress >	imported_files;
		imported_files.AddOrReplace( FileAddress( filename ) );

		CHECK_RETURN( _RecursiveParseCLFile( filename, src, imported_files ) );


		// create opencl program
		src.Insert( string("#line 1 \"") << FileAddressUtils::GetNameAndExt( filename ) << "\"\n", 0 );

		if ( not preprocessor.Empty() ) {
			src.Insert( preprocessor, 0 );
			src.Insert( "#line 1 \"preprocessor\"\n", 0 );
		}
		src.Insert( _clHeader, 0 );

		if ( not options.Empty() ) {
			src.Insert( "#line 1 \"options\"\n", 0 );
			src.Insert( options, 0 );
		}

		CHECK_RETURN( program->Create( src.cstr() ) );

		const bool builded = program->Build( _clOptions );

		//if ( not builded )
		{
			string	log;
			log << "In file \"" << filename << "\":\n";

			const usize min_log_size = log.Length() + 3;

			CHECK_RETURN( program->GetBuildLog( log ) );

			if ( not builded ) {
				RETURN_ERROR( log );
			}
			else 
			if ( log.Length() > min_log_size ) {
				WARNING( log.cstr() );
			}
		}

		return true;
	}
	
/*
=================================================
	_ReadFile
=================================================
*/
	bool ShaderManager::_ReadFile (uni_c_string filename, OUT string &buf)
	{
		RFilePtr	file;
		CHECK_VALUE( (file = File::HddRFile::New( filename )).IsNotNull() );
		
		const usize	len = file->Size();

		buf.Reserve( len + 1 );
		CHECK_RETURN( file->Read( buf.ptr(), len ) );
		buf.SetLength( len );

		return true;
	}
	
/*
=================================================
	_RecursiveParseCLFile
----
	Warning: #include work in commented code too!
=================================================
*/
	bool ShaderManager::_RecursiveParseCLFile (uni_c_string filename, OUT string &src, INOUT set< FileAddress > &importedFiles)
	{
		static const uni_c_string	s_import( "#include" );
		
		const uni_c_string			path  = FileAddressUtils::GetPath( filename );
		const uni_c_string			fname = FileAddressUtils::GetNameAndExt( filename );

		CHECK_RETURN( _ReadFile( filename, src ) );
		

		// find #import directives
		array< svec2 >	imports;	// position in source and line number
		imports.Reserve( 128 );

		for (usize pos = 0; src.Find( s_import, pos, pos ); ++pos)
		{
			const usize	prev_pos	= imports.Empty() ? 0 : imports.Back().x;
			const usize	prev_lines	= imports.Empty() ? 1 : imports.Back().y;

			const usize lines		= StringParser::CalculateNumberOfLines(
										uni_c_string( src.cstr() + prev_pos, pos - prev_pos) );

			imports.PushBack( svec2( pos, lines + prev_lines ) );
		}


		// insert line numbers
		usize	offset = 0;
		string	buf;
		buf.Reserve( 1<<10 );

		FOR( i, imports )
		{
			const usize	begin	= imports[i].x + offset;
			const usize	line	= imports[i].y;
			usize		pos		= begin;
			usize		index	= 0;

			pos += s_import.Length();
			
			uni_c_string	import_lib;
			CHECK_RETURN( _ParseString( src, pos, import_lib ) );
			
			StringParser::ToNextLine( src, pos );
			
			bool			is_imported	= ( import_lib == "cl.h" );

			FileAddress		import_lib_address( import_lib );
			import_lib_address.AddBasePath( path );

			if ( not is_imported ) {
				is_imported = importedFiles.FindIndex( import_lib_address, index );
			}

			// remove #include line
			src.Erase( begin, pos - begin );

			offset -= (pos - begin);
			
			// is file was already imported?
			if ( not is_imported )
			{
				// import source
				buf.ClearMem();
				CHECK_RETURN( _RecursiveParseCLFile( import_lib_address.Full(), buf, importedFiles ) );
				CHECK_VALUE( not buf.Empty() );

				buf.Insert( string("#line 1 \"") << import_lib_address.NameWithExt() << "\"\n", 0 );
				buf.Append( string("#line ") << (line+1) << " \"" << fname << "\"\n" );

				src.Insert( buf, begin );
				
				importedFiles.AddOrReplace( import_lib_address );

				offset += buf.Length();
			}
		}

		return true;
	}

/*
=================================================
	_BuildGLSLProprocessor
=================================================
*/
	void ShaderManager::_BuildGLSLProprocessor (OUT string &str)
	{
		str.Reserve( 1024 );
		str <<
			"#version " << GLSYSTEM()->GLSLVersion() << " core\n"
			"#define GPU_MODEL_UNKNOWN	0\n"
			"#define GPU_MODEL_AMD		1\n"
			"#define GPU_MODEL_NVIDIA	2\n"
			"#define GPU_MODEL_INTEL	3\n"
			"#define GLSL_VERSION		__VERSION__\n";

		const e_gpu_model::type	gpu_model = GLSYSTEM()->GetGPUModel();

		switch ( gpu_model )
		{
			case e_gpu_model::AMD :
				str <<
					"#define GPU_MODEL		GPU_MODEL_AMD\n"
					"#define static_const	const\n";
				break;

			case e_gpu_model::NVIDIA :
				str <<
					"#define GPU_MODEL		GPU_MODEL_NVIDIA\n"
					"#define static_const	uniform\n";			// it work faster then "const"
				break;

			case e_gpu_model::INTEL :
				str <<
					"#define GPU_MODEL		GPU_MODEL_INTEL\n"
					"#define static_const	const\n";
				break;

			default :
				str <<
					"#define GPU_MODEL		GPU_MODEL_UNKNOWN\n"
					"#define static_const	const\n";
				break;
		}
		/*
		str <<
			"#define not	!\n"
			"#define and	&&\n"
			"#define or		||\n";
		*/
		str <<
			"\n/* compiler options */\n";

		DEBUG_ONLY(
			str <<
				"#pragma optimize(off)\n"
				"#pragma debug(on)\n";
		);
		RELEASE_ONLY(
			str <<
				"#pragma optimize(on)\n"
				"#pragma debug(off)\n"
				"#if GPU_MODEL == GPU_MODEL_NVIDIA\n"
				"#	pragma optionNV(fastmath on)\n"
				"#	pragma optionNV(fastprecision on)\n"
				"#	pragma optionNV(ifcvt none)\n"
				"#	pragma optionNV(inline all)\n"
				"#	pragma optionNV(strict on)\n"
				"#	pragma optionNV(unroll all)\n"
				"#endif\n";
		);

		str << "/* end of build-in header */\n\n";
	}
	
/*
=================================================
	_BuildCLPreprocessor
=================================================
*/
	void ShaderManager::_BuildCLPreprocessor (OUT string &str, OUT cl_prog_options::type &opt)
	{
		const bool image_3d_supported = false; //CLDEVICE()->IsExtSupported( "cl_khr_3d_image_writes" );

		str.Reserve( 1024 );
		
		str << "#line 1 \"build-in\"\n";
		
		// extensions //
		if ( image_3d_supported ) {
			str <<	"#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n\n";
		}

		// temporary macroses //
		// FP
		switch ( sizeof(real) )
		{
			case sizeof(half) :
				str << "#if !defined(FP)\n#define FP 16\n#endif\n\n";
				break;

			case sizeof(double) :
				str << "#if !defined(FP)\n#define FP 64\n#endif\n\n";
				break;

			case sizeof(float) :
			default :
				str << "#if !defined(FP)\n#define FP 32\n#endif\n\n";
				break;
		}

		// Check Image Format
		str <<
			"#define __IsFloatImage( image )		__IsFloatImageFmt( get_image_channel_data_type( image ) ) \n"
			"#define __IsSignedIntImage( image )	__IsSignedIntImageFmt( get_image_channel_data_type( image ) ) \n"
			"#define __IsUnsignedIntImage( image )	__IsUnsignedIntImageFmt( get_image_channel_data_type( image ) ) \n"
			"#define __ImageChannelsCount( image )	__ImageChannelsCountFromChannelOrder( get_image_channel_order( image ) ) \n"
			
			"bool __IsFloatImageFmt (int fmt) {"
			"	return	(fmt == CLK_SNORM_INT8) | (fmt == CLK_SNORM_INT16) |"
			"			(fmt == CLK_UNORM_INT8) | (fmt == CLK_UNORM_INT16) |"
			"			(fmt == CLK_UNORM_SHORT_565) | (fmt == CLK_UNORM_SHORT_555) |"
			"			(fmt == CLK_UNORM_INT_101010) | (fmt == CLK_HALF_FLOAT) |"
			"			(fmt == CLK_FLOAT);"
			"}\n"
			"bool __IsSignedIntImageFmt (int fmt) {"
			"	return (fmt == CLK_SIGNED_INT8) | (fmt == CLK_SIGNED_INT16) | (fmt == CLK_SIGNED_INT32);"
			"}\n"
			"bool __IsUnsignedIntImageFmt (int fmt) {"
			"	return (fmt == CLK_UNSIGNED_INT8) | (fmt == CLK_UNSIGNED_INT16) | (fmt == CLK_UNSIGNED_INT32);"
			"}\n"
			"int __ImageChannelsCountFromChannelOrder (int channelOrder) {"
			"	switch( channelOrder ) {"
			"		case CLK_A :"
			"		case CLK_R :"
			"		case CLK_Rx :"
			"			return 1;"
			"		case CLK_RG :"
			"		case CLK_RGx :"
			"		case CLK_RA :"
			"			return 2;"
			"		case CLK_RGB :"
			"		case CLK_RGBx :"
			"			return 3;"
			"		case CLK_RGBA :"
			"		case CLK_ARGB :"
			"		case CLK_BGRA :"
			"			return 4;"
			"		case CLK_INTENSITY :"
			"		case CLK_LUMINANCE :"
			"			return 1;"
			"		default :"
			"			return 0;"
			"	}"
			"}\n"
			
			"#if DEBUG == 1\n"
			"#	define __IMGNOISE( t )"
			"		(t)( sin( (float)(get_global_id(0) + get_global_offset(0)) ) ) \n"

			"#	define safe_write_imagef( image, coord, color ) { "
			"		if ( __IsFloatImage( (image) ) )		{ write_imagef( (image), (coord), (color) ); } else "
			"		if ( __IsSignedIntImage( (image) ) )	{ write_imagei( (image), (coord), __IMGNOISE(int4) ); } else "
			"		if ( __IsUnsignedIntImage( (image) ) )	{ write_imageui( (image), (coord), __IMGNOISE(uint4) ); } "
			"	}\n"

			"#	define safe_write_imagei( image, coord, color ) { "
			"		if ( __IsSignedIntImage( (image) ) )	{ write_imagei( (image), (coord), (color) ); } else "
			"		if ( __IsFloatImage( (image) ) )		{ write_imagef( (image), (coord), __IMGNOISE(float4) ); } else "
			"		if ( __IsUnsignedIntImage( (image) ) )	{ write_imageui( (image), (coord), __IMGNOISE(uint4) ); } "
			"	}\n"
			
			"#	define safe_write_imageui( image, coord, color ) { "
			"		if ( __IsUnsignedIntImage( (image) ) )	{ write_imageui( (image), (coord), (color) ); } else "
			"		if ( __IsFloatImage( (image) ) )		{ write_imagef( (image), (coord), __IMGNOISE(float4) ); } else "
			"		if ( __IsSignedIntImage( (image) ) )	{ write_imagei( (image), (coord), __IMGNOISE(int4) ); } "
			"	}\n"

			"#	define safe_write_imaged( image, coord, color ) { "
			"		int channels = __ImageChannelsCount( (image) ); "
			"		if ( (channels == 2 or channels == 4) and __IsSignedIntImage( (image) ) )	{ write_imaged( (image), (coord), (color) ); } else "
			"		if ( __IsSignedIntImage( (image) ) )	{ write_imagei( (image), (coord), __IMGNOISE(int4) ); } else "
			"		if ( __IsFloatImage( (image) ) )		{ write_imagef( (image), (coord), __IMGNOISE(float4) ); } else "
			"		if ( __IsSignedIntImage( (image) ) )	{ write_imagei( (image), (coord), __IMGNOISE(int4) ); } "
			"	}\n"

			"#else\n"
			
			"#	define safe_write_imagef( image, coord, color )		write_imagef( (image), (coord), (color) ) \n"
			"#	define safe_write_imagei( image, coord, color )		write_imagei( (image), (coord), (color) ) \n"
			"#	define safe_write_imageui( image, coord, color )	write_imageui( (image), (coord), (color) ) \n"
			"#	define safe_write_imaged( image, coord, color )		write_imaged( (image), (coord), (color) ) \n"

			"#endif\n\n";

		
		// Real type //
		str	<<
			"#if FP == 16\n"
			"# pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
			"# define REAL_IS_HALF		1\n"
			"# define real				half\n"
			"# define real2				half2\n"
			"# define real3				half3\n"
			"# define real4				half4\n"
			"# define real8				half8\n"
			"# define real16			half16\n"
			"# define convert_real		convert_half\n"
			"# define convert_real2		convert_half2\n"
			"# define convert_real3		convert_half3\n"
			"# define convert_real4		convert_half4\n"
			"# define convert_real8		convert_half8\n"
			"# define convert_real16	convert_half16\n"
			"# define as_real			as_half\n"
			"# define as_real2			as_half2\n"
			"# define as_real3			as_half3\n"
			"# define as_real4			as_half4\n"
			"# define as_real8			as_half8\n"
			"# define as_real16			as_half16\n"
			"# define write_image(image, coord, color)		write_imagef( (image), (coord), convert_float4(color) )\n"
			"# define read_image(image, sampler, coord)		convet_half4( read_imagef( (image), (sampler), (coord) ) )\n"
			"#endif\n\n"

			"#if FP == 64\n"
			"# pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
			"# define REAL_IS_DOUBLE	1\n"
			"# define real				double\n"
			"# define real2				double2\n"
			"# define real3				double3\n"
			"# define real4				double4\n"
			"# define real8				double8\n"
			"# define real16			double16\n"
			"# define convert_real		convert_double\n"
			"# define convert_real2		convert_double2\n"
			"# define convert_real3		convert_double3\n"
			"# define convert_real4		convert_double4\n"
			"# define convert_real8		convert_double8\n"
			"# define convert_real16	convert_double16\n"
			"# define as_real			as_double\n"
			"# define as_real2			as_double2\n"
			"# define as_real3			as_double3\n"
			"# define as_real4			as_double4\n"
			"# define as_real8			as_double8\n"
			"# define as_real16			as_double16\n"
			"# define write_image		write_imaged\n"
			"# define read_image		read_imaged\n"
			"#endif\n\n"

			"#if FP == 32\n"
			"# define REAL_IS_FLOAT		1\n"
			"# define real				float\n"
			"# define real2				float2\n"
			"# define real3				float3\n"
			"# define real4				float4\n"
			"# define real8				float8\n"
			"# define real16			float16\n"
			"# define convert_real		convert_float\n"
			"# define convert_real2		convert_float2\n"
			"# define convert_real3		convert_float3\n"
			"# define convert_real4		convert_float4\n"
			"# define convert_real8		convert_float8\n"
			"# define convert_real16	convert_float16\n"
			"# define as_real			as_float\n"
			"# define as_real2			as_float2\n"
			"# define as_real3			as_float3\n"
			"# define as_real4			as_float4\n"
			"# define as_real8			as_float8\n"
			"# define as_real16			as_float16\n"
			"# define write_image		write_imagef\n"
			"# define read_image		read_imagef\n"
			"#endif\n\n";


		// Undef temporary macroses //
		str <<
			"#undef FP \n\n";

		// read/write double to image
		str << 
			" union __Int4ToDouble2 { int4 i; double2 d; };\n"

			"# define write_imaged( image, coord, color) {\\\n"
			"	union __Int4ToDouble2 u;\\\n"
			"	u.d = color;\\\n"
			"	write_imagei( image, coord, u.i );\\\n"
			" }\n"
			
			"# define read_imaged( image, sampler, coord) {\\\n"
			"	union __Int4ToDouble2 u;\\\n"
			"	u.i = read_imagei( image, sampler, coord );\\\n"
			"	return u.d;\\\n"
			" }\n"
			;

		str <<
			"#define OUT \n"
			"#define INOUT \n"
			"#define ASSERT(x) \n";

		str << "/* end of build-in header */\n\n";

		opt  = cl_prog_options::CL_1_1;
		opt |= cl_prog_options::WARNING_AS_ERROR;

		DEBUG_ONLY(
			//opt |= cl_prog_options::OPT_DISABLE;
		)
		RELEASE_ONLY(
			opt |= cl_prog_options::OPT_MAD_ENABLE;
		)
	}



//------------------------------- ShaderLoader ------------------------------//
/*
=================================================
	Load
=================================================
*/
	bool ShaderLoader::Load (GLProgramPtr &program, uni_c_string filename, uni_c_string preprocessor)
	{
		return ShaderManager::Instance()->Load( program, filename, preprocessor );
	}
	
/*
=================================================
	Load
=================================================
*/
	bool ShaderLoader::Load (CLProgramPtr &program, uni_c_string filename, uni_c_string preprocessor, uni_c_string options)
	{
		return ShaderManager::Instance()->Load( program, filename, preprocessor, options );
	}

/*
=================================================
	FreeCache
=================================================
*/
	void ShaderLoader::FreeCache ()
	{
		return ShaderManager::Instance()->FreeCache();
	}


}	// Framework

#include <logger/ux_logger_lib.h>
